Add Cooperative Groups double buffered tile example#486
Open
j-stephan wants to merge 1 commit into
Open
Conversation
Signed-off-by: Jan Stephan <jan.stephan@amd.com>
randyh62
approved these changes
Jul 17, 2026
Comment on lines
+10
to
+12
| The async load of the next tile is issued into the other buffer while the current tile is consumed, | ||
| and a split barrier separates the moment a thread has finished reading a buffer from the moment the | ||
| block guarantees that every thread is done. The kernel applies the element-wise operation |
Contributor
There was a problem hiding this comment.
Suggested change
| The async load of the next tile is issued into the other buffer while the current tile is consumed, | |
| and a split barrier separates the moment a thread has finished reading a buffer from the moment the | |
| block guarantees that every thread is done. The kernel applies the element-wise operation | |
| The async load of the next tile is issued into the second buffer while the current tile is consumed. | |
| A split barrier separates the moment a thread has finished reading a buffer from the moment the | |
| block guarantees that every thread is done. The kernel applies the element-wise operation |
Contributor
There was a problem hiding this comment.
The second sentence could be a little more clear as to the split barrier separating the two moments: the moment the thread reads the buffer, and the moment all threads have completed the operation.
| ## Description | ||
|
|
||
| This program showcases a double-buffered tile load pipeline built from two cooperative groups | ||
| APIs: the group-collective `cooperative_groups::memcpy_async` and the split barrier |
Contributor
There was a problem hiding this comment.
is the term spilt barrier well understood? Should it be defined here at the start?
Comment on lines
+26
to
+28
| This example targets the AMD/HIP (ROCm) backend: the NVIDIA path of | ||
| `<hip/cooperative_groups/memcpy_async.h>` is not yet implemented (the header carries a TODO for it), | ||
| and it requires a ROCm version recent enough to ship that public header. |
Contributor
There was a problem hiding this comment.
Suggested change
| This example targets the AMD/HIP (ROCm) backend: the NVIDIA path of | |
| `<hip/cooperative_groups/memcpy_async.h>` is not yet implemented (the header carries a TODO for it), | |
| and it requires a ROCm version recent enough to ship that public header. | |
| This example targets the AMD/HIP (ROCm) backend, | |
| and it requires a ROCm version recent enough to ship that public header. |
|
|
||
| ### Application flow | ||
|
|
||
| 1. A number of compile-time constants define the tile size (also the block size), the number of |
Contributor
There was a problem hiding this comment.
Suggested change
| 1. A number of compile-time constants define the tile size (also the block size), the number of | |
| 1. A number of compile-time constants define the tile size and block size, the number of |
Comment on lines
+55
to
+58
| sizeof(float)`). The copy is asynchronous and HIP exposes no separate wait handle, so its | ||
| completion is enforced by a following group barrier (`block.sync()` or the split barrier's | ||
| `barrier_wait`). On hardware or compilers without the asynchronous LDS builtins it falls back to a | ||
| correct traditional per-thread copy, so it always produces correct results. |
Contributor
There was a problem hiding this comment.
Suggested change
| sizeof(float)`). The copy is asynchronous and HIP exposes no separate wait handle, so its | |
| completion is enforced by a following group barrier (`block.sync()` or the split barrier's | |
| `barrier_wait`). On hardware or compilers without the asynchronous LDS builtins it falls back to a | |
| correct traditional per-thread copy, so it always produces correct results. | |
| sizeof(float)`). The copy is asynchronous and HIP exposes no separate wait handle, so the copy | |
| completion is enforced by a following group barrier (`block.sync()` or `barrier_wait`). | |
| On hardware or compilers without the asynchronous LDS builtins it falls back to a | |
| correct traditional per-thread copy, so it always produces correct results. |
Contributor
There was a problem hiding this comment.
It is not clear to me what "it" refers to in "it falls back to"
| completion is enforced by a following group barrier (`block.sync()` or the split barrier's | ||
| `barrier_wait`). On hardware or compilers without the asynchronous LDS builtins it falls back to a | ||
| correct traditional per-thread copy, so it always produces correct results. | ||
| - The split barrier decomposes a block barrier into two phases. `thread_block::barrier_arrive` |
Contributor
There was a problem hiding this comment.
This paragraph (or bullet) would be very helpful at the start of this topic for those who might be learning.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
In ROCm 7.14, we introduced new Cooperative Groups features. This example demonstrates how to use
memcpy_async.Technical Details
Not supported for HIP SDK or CUDA.
Test Plan
Tested locally on gfx1100 lab machine.
Test Result
Works.
Added/Updated documentation?
Included Visual Studio files?
Submission Checklist